#define PRECISION double

#ifdef cl_amd_fp64
#pragma OPENCL EXTENSION cl_amd_fp64 : enable
#else
#pragma OPENCL EXTENSION cl_khr_fp64 : enable
#endif




__kernel void OpenCLTemplateKernel(__global  PRECISION * Input,
      					__global PRECISION * Filter,
      					__global  PRECISION * Output,
      					const int nInputWidth,
      					const int nFilterWidth) 
{
   const int nWidth = get_global_size(0);
   const int xOut = get_global_id(0);
   const int yOut = get_global_id(1);
   const int xInTopLeft = xOut;
   const int yInTopLeft = yOut;
   PRECISION sum = 0;
    for (int r = 0; r < nFilterWidth; r++)
    {
  	  //since linear array of floats, need to index correct element
  	  const int idxFtmp = r * nFilterWidth;
  	  const int yIn = yInTopLeft + r;
  	  const int idxIntmp = yIn * nInputWidth + xInTopLeft;
  	  for (int c = 0; c < nFilterWidth; c++)
  	  {
  	  	const int idxF  = idxFtmp  + c;
  	  	const int idxIn = idxIntmp + c;
  	  	sum += Filter[xOut*idxF+c]*Input[xOut*idxIn+c];

  	  }
   }
   const int idxOut = yOut * nWidth + xOut;
   Output[yOut*idxOut+xOut] = sum;
}